#include "StdAfx.h"
#include "OpenCL.h"
#include "FileIO.h"
#include <iostream>



//-------------------------------------------------------------------------------------------------
OpenCL::OpenCL(cl_device_type type, cl_command_queue_properties prop)
:	error						(0),
	num_platforms			(0),
	num_devices				(0),
	context					(NULL),
	queue						(NULL),
	program					(0),
	kernel					(0),
	mem						(0),
	totalMem					(0)
{
	error		= clGetPlatformIDs(8, platform, &num_platforms);

	error		= clGetDeviceIDs(platform[0], type, 1, &device[0], &num_devices);

	context	= clCreateContext(0, 1, device, NULL, NULL, &error);

	queue		= clCreateCommandQueue(context, device[0], prop, &error);
}


//-------------------------------------------------------------------------------------------------
OpenCL::~OpenCL()
{
	for(unsigned i = 0; i < kernel.size(); i++)
		error = clReleaseKernel(kernel[i]);

	for(unsigned i = 0; i < program.size(); i++)
		error = clReleaseProgram(program[i]);

	for(unsigned i = 0; i < mem.size(); i++)
		error = clReleaseMemObject(mem[i]);

	error = clReleaseCommandQueue(queue);
	error = clReleaseContext(context);
}


//-------------------------------------------------------------------------------------------------
void OpenCL::programFromSource(const char *filename, int i)
{
	if(i >= program.size())
		program.resize(i+8);
	else if(mem[i])
		error = clReleaseProgram(program[i]);				// release old program

// Create OpenCL program with source code
	std::vector<char> sourceFile;
	if(!FileLoad(filename, sourceFile))
		return;

	size_t size = sourceFile.size();
	const char *source = &sourceFile[0];
	program[i] = clCreateProgramWithSource(context, 1, &source, &size, &error);
}

//-------------------------------------------------------------------------------------------------
bool OpenCL::programFromBinary(const char *filename, int i)
{
	if(i >= program.size())
		program.resize(i+8);
	else if(mem[i])
		error = clReleaseProgram(program[i]);				// release old program

	std::vector<unsigned char> file;
	if(!FileLoad(filename, file))
		return false;

	size_t size = file.size();
	const unsigned char *data = &file[0];

	cl_int status;
	program[i] = clCreateProgramWithBinary(context, 1, &device[0], &size, &data, &status, &error);

	return status == CL_SUCCESS;
}

//-------------------------------------------------------------------------------------------------
void OpenCL::programSaveBinary(const char *filename, int i)
{
	if(i >= program.size())
		return;

	size_t size[8];

	error = clGetProgramInfo(program[i], CL_PROGRAM_BINARY_SIZES, sizeof(size), &size, NULL);

	std::vector<unsigned char> data(size[0]);
	unsigned char *ptr = &data[0];

	error = clGetProgramInfo(program[i], CL_PROGRAM_BINARIES, sizeof(ptr), &ptr, NULL);

	FileSave(filename, data);
}

//-------------------------------------------------------------------------------------------------
void OpenCL::programBuild(int i, const char *options)
{
// Compile the program (OpenCL JIT compilation)
	error = clBuildProgram(program[i], 1, &device[0], options, NULL, NULL);
}

//-------------------------------------------------------------------------------------------------
bool OpenCL::kernelCreate(int i, const char *name, int prg)
{
	if(name == NULL)
	{
		int n = kernel.size();

		kernel.resize(n+64);

		cl_uint num_kernels_ret;
		cl_int result = clCreateKernelsInProgram(program[prg], 64, &kernel[n], &num_kernels_ret);
		return result == CL_SUCCESS;
	}
	else
	{
		if(i >= kernel.size())
			kernel.resize(i+8);
		else if(kernel[i])
			error = clReleaseKernel(kernel[i]);					// release old kernel

	// Assemble the kernel
		kernel[i] = clCreateKernel(program[prg], name, &error);
		return error == 0;
	}
}


//-------------------------------------------------------------------------------------------------
bool OpenCL::kernelSetArg(int i, cl_uint argIndex, size_t argSize, const void *argValue)
{
	cl_int result = clSetKernelArg(kernel[i], argIndex, argSize, argValue);
	return result == CL_SUCCESS;
}

//-------------------------------------------------------------------------------------------------
bool OpenCL::kernelSetArgMem(int i, cl_uint argIndex, int memIndex)
{
	cl_int result = clSetKernelArg(kernel[i], argIndex, sizeof(cl_mem), &mem[memIndex]);
	return result == CL_SUCCESS;
}

//-------------------------------------------------------------------------------------------------
bool OpenCL::kernelEnqueue1D(int i, size_t global_work_size, size_t local_work_size, cl_event *e)
{
	size_t *ptr = local_work_size == 0 ? NULL : &local_work_size;
	cl_int result = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, &global_work_size, ptr, 0, NULL, e);
	return result == CL_SUCCESS;
}

//-------------------------------------------------------------------------------------------------
int OpenCL::kernelWorkGroupSize(int i)
{
	int ret;
	cl_int result = clGetKernelWorkGroupInfo(kernel[i], device[0], CL_KERNEL_WORK_GROUP_SIZE, sizeof(ret), &ret, NULL);
	return ret;
}



//-------------------------------------------------------------------------------------------------
bool OpenCL::bufferAllocate	(int i, size_t size, cl_mem_flags flags)
{
	if(i >= mem.size())
		mem.resize(i+8);
	else if(mem[i])
		error = clReleaseMemObject(mem[i]);					// release old buffer

// Allocate output memory on GPU
	mem[i] = clCreateBuffer		(context, flags, size, NULL, &error);

	if(error==0)
	{
		totalMem += size;
		return true;
	}
	else
	{
		std::cout << "OpenCL::bufferAllocate("<<(size/(1024*1024))<<" MiB) failed! (totalMem = "<<totalMem<<")" << std::endl;
		return false;
	}
}

//-------------------------------------------------------------------------------------------------
bool OpenCL::bufferAllocate	(int i, void *src, size_t size, cl_mem_flags flags)
{
	if(i >= mem.size())
		mem.resize(i+8);
	else if(mem[i])
		error = clReleaseMemObject(mem[i]);					// release old buffer

// Allocate GPU memory for source vectors AND initialize from CPU memory
	mem[i] = clCreateBuffer		(context, flags, size, src, &error);

	if(error==0)
	{
		totalMem += size;
		return true;
	}
	else
	{
		std::cout << "OpenCL::bufferAllocate("<<(size/(1024*1024))<<" MiB) failed! (totalMem = "<<totalMem<<")" << std::endl;
		return false;
	}
}

//-------------------------------------------------------------------------------------------------
void OpenCL::bufferRead			(int i, void *dst, size_t size, cl_bool block, cl_event *e)
{
	error = clEnqueueReadBuffer(queue, mem[i], block, 0, size, dst, 0, NULL, e);
}

//-------------------------------------------------------------------------------------------------
void OpenCL::bufferWrite		(int i, void *src, size_t size, cl_bool block, cl_event *e)
{
	error = clEnqueueWriteBuffer(queue, mem[i], block, 0, size, src, 0, NULL, e);
}

//-------------------------------------------------------------------------------------------------
void* OpenCL::bufferMap			(int i, cl_map_flags flags, cl_bool block, cl_event *e)
{
	return clEnqueueMapBuffer	(queue, mem[i], block, flags, 0, bufferInfo(i), 0, NULL, e, &error);
}

//-------------------------------------------------------------------------------------------------
void OpenCL::bufferUnmap		(int i, void *ptr, cl_event *e)
{
	error = clEnqueueUnmapMemObject(queue, mem[i], ptr, 0, NULL, e);
}

//-------------------------------------------------------------------------------------------------
size_t OpenCL::bufferInfo		(int i, cl_mem_info param_name)
{
	size_t ret;
	error = clGetMemObjectInfo (mem[i], param_name, sizeof(ret), &ret, NULL);
	return ret;
}

//-------------------------------------------------------------------------------------------------
void OpenCL::bufferRelease		(int i)
{
	if(i < mem.size() && mem[i])
	{
		size_t size = bufferInfo(i);		// get buffer size
		error = clReleaseMemObject(mem[i]);
		if(error==0)
			totalMem -= size;			
	}
}


//-------------------------------------------------------------------------------------------------
bool OpenCL::flush()
{
	cl_int result = clFlush(queue);
	return result == CL_SUCCESS;
}


//-------------------------------------------------------------------------------------------------
bool OpenCL::finish()
{
	cl_int result = clFinish(queue);
	return result == CL_SUCCESS;
}







//-------------------------------------------------------------------------------------------------
cl_int OpenCL::EventWait(cl_event e)
{
	return clWaitForEvents(1, &e);
}


//-------------------------------------------------------------------------------------------------
cl_int OpenCL::EventInfo(cl_event e, cl_event_info nfo)
{
	cl_int result;
	cl_int error = clGetEventInfo(e, nfo, sizeof(result), &result, NULL);
	return result;
}




//-------------------------------------------------------------------------------------------------
cl_ulong OpenCL::EventStartTime	(cl_event e)
{
	cl_ulong t;
	cl_int error = clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_START, sizeof(t), &t, NULL);
	return t;
}

//-------------------------------------------------------------------------------------------------
cl_ulong OpenCL::EventEndTime		(cl_event e)
{
	cl_ulong t;
	cl_int error = clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_END, sizeof(t), &t, NULL);
	return t;
}



//-------------------------------------------------------------------------------------------------
void OpenCL::PrintDeviceName(cl_device_id device)
{
	cl_int err;

	char *buffer = new char[1024];

	err = clGetDeviceInfo(device,CL_DEVICE_NAME,1024,buffer,NULL);
	std::cout << "NAME = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_VENDOR,1024,buffer,NULL);
	std::cout << "VENDOR = " << buffer << std::endl;

	delete[] buffer;
}

//-------------------------------------------------------------------------------------------------
void OpenCL::PrintDeviceInfo(cl_device_id device)
{
	cl_int err;

	char *buffer = new char[1024];
	err = clGetDeviceInfo(device,CL_DEVICE_NAME,1024,buffer,NULL);
	std::cout << "NAME = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_VENDOR,1024,buffer,NULL);
	std::cout << "VENDOR = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_VERSION,1024,buffer,NULL);
	std::cout << "VERSION = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DRIVER_VERSION,1024,buffer,NULL);
	std::cout << "DRIVER_VERSION = " << buffer << std::endl;

	cl_bool b;
	err = clGetDeviceInfo(device,CL_DEVICE_ENDIAN_LITTLE,sizeof(b),&b,NULL);
	std::cout << "ENDIAN_LITTLE = " << b << std::endl;

	cl_bitfield cap;
	err = clGetDeviceInfo(device,CL_DEVICE_EXECUTION_CAPABILITIES,sizeof(cap),&cap,NULL);
	std::cout << "EXECUTION_CAPABILITIES: NATIVE_KERNEL = " << (cap & CL_EXEC_NATIVE_KERNEL) << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_EXTENSIONS,1024,buffer,NULL);
	std::cout << "EXTENSIONS = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PROFILE,1024,buffer,NULL);
	std::cout << "PROFILE = " << buffer << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_TYPE,sizeof(cap),&cap,NULL);
	std::cout << "TYPE: CPU = " << (CL_DEVICE_TYPE_CPU&cap) << ", GPU = " << (CL_DEVICE_TYPE_GPU&cap) << ", ACCELERATOR = " << (CL_DEVICE_TYPE_ACCELERATOR&cap) << ", DEFAULT = " << (CL_DEVICE_TYPE_DEFAULT&cap) << std::endl;

	cl_uint ui;
	err = clGetDeviceInfo(device,CL_DEVICE_ADDRESS_BITS,sizeof(ui),&ui,NULL);
	std::cout << "ADDRESS_BITS = " << ui << std::endl;

	cl_ulong ul;
	err = clGetDeviceInfo(device,CL_DEVICE_GLOBAL_MEM_CACHE_SIZE,sizeof(ul),&ul,NULL);
	std::cout << "GLOBAL_MEM_CACHE_SIZE = " << ul << std::endl;

	cl_device_mem_cache_type type;
	err = clGetDeviceInfo(device,CL_DEVICE_GLOBAL_MEM_CACHE_TYPE,sizeof(type),&type,NULL);
	std::cout << "GLOBAL_MEM_CACHE_TYPE: NONE = " << (CL_NONE&type) << ", READ_ONLY = " << (CL_READ_ONLY_CACHE&type) << ", READ_WRITE = " << (CL_READ_WRITE_CACHE&type) << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE,sizeof(ui),&ui,NULL);
	std::cout << "GLOBAL_MEM_CACHELINE_SIZE = " << ui << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(ul),&ul,NULL);
	std::cout << "GLOBAL_MEM_SIZE = " << ul << " = " << (ul/(1024*1024)) << " MiB" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_MAX_MEM_ALLOC_SIZE,sizeof(ul),&ul,NULL);
	std::cout << "MAX_MEM_ALLOC_SIZE = " << ul << " = " << (ul/(1024*1024)) << " MiB" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_LOCAL_MEM_SIZE,sizeof(ul),&ul,NULL);
	std::cout << "LOCAL_MEM_SIZE = " << ul << " = " << (ul/(1024)) << " KiB" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_LOCAL_MEM_TYPE,sizeof(type),&type,NULL);
	std::cout << "LOCAL_MEM_TYPE: LOCAL = " << (CL_LOCAL&type) << ", GLOBAL = " << (CL_GLOBAL&type) << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_MAX_CLOCK_FREQUENCY,sizeof(ui),&ui,NULL);
	std::cout << "MAX_CLOCK_FREQUENCY = " << ui << " MHz" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,sizeof(ul),&ul,NULL);
	std::cout << "MAX_CONSTANT_BUFFER_SIZE = " << ul  << " = " << (ul/(1024)) << " KiB" << std::endl;

	size_t size;
	err = clGetDeviceInfo(device,CL_DEVICE_MAX_PARAMETER_SIZE,sizeof(size),&size,NULL);
	std::cout << "MAX_PARAMETER_SIZE = " << size  << " = " << (size/(1024)) << " KiB" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(ui),&ui,NULL);
	std::cout << "MAX_WORK_ITEM_DIMENSIONS = " << ui  << std::endl;

	size_t sizes[3];
	err = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(sizes),&sizes,NULL);
	std::cout << "MAX_WORK_ITEM_SIZES = [" << sizes[0] << " | " << sizes[1] << " | " << sizes[2] << "]" << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_CHAR = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_SHORT = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_INT = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_LONG = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_FLOAT = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE,sizeof(ui),&ui,NULL);
	std::cout << "PREFERRED_VECTOR_WIDTH_DOUBLE = " << ui  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_PROFILING_TIMER_RESOLUTION,sizeof(size),&size,NULL);
	std::cout << "PROFILING_TIMER_RESOLUTION = " << size << "ns"  << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_QUEUE_PROPERTIES,sizeof(cap),&cap,NULL);
	std::cout << "QUEUE_PROPERTIES = OUT_OF_ORDER_EXEC_MODE_ENABLE = " << (cap&CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) << std::endl;

	err = clGetDeviceInfo(device,CL_DEVICE_SINGLE_FP_CONFIG,sizeof(cap),&cap,NULL);
	std::cout << "SINGLE_FP_CONFIG: DENORM = " << (CL_FP_DENORM&cap) << ", ROUND_TO_ZERO = " << (CL_FP_ROUND_TO_ZERO&cap) << ", ROUND_TO_INF = " << (CL_FP_ROUND_TO_INF&type) << ", FMA = " << (CL_FP_FMA&type) << std::endl;

//----
	std::cout << std::endl;
	delete[] buffer;
}
